[Doc][Polish] gemm optimize by 2d thread tile #56
Conversation
muyuuuu
force-pushed
the
gemm_tiled_2d
branch
from
cb694d7 to
e272187
Compare
docs/07_optimize_matmul/README.md
Outdated
| @@ -88,7 +88,7 @@ nvcc -o matmul_shared matmul_shared.cu | |||
|
|
|||
|  | |||
|
|
|||
| 上图中，A 和 B 都是 7x7 的矩阵。当每一个线程只计算一个结果的时候，我们需要从 A 中读取 7 个数据，从 B 中读取 7 个数据，从 C 中读取 1 个数据，然后写一次 C。这样的话，每个线程需要读取 15 个数据，写一次数据。如果我们每一个线程计算 4 个结果，那么我们需要从 A 中读取 14 个数据，从 B 中读取 14 个数据，从 C 中读取 4 个数据，然后写 4 次 C。这样的话，每个线程需要读取 32 个数据，写 4 次数据。计算每个线程的平方结果比计算结果的列更有效，因为这样我们可以共享更多的输入。 | |||
| 上图中，A 和 B 都是 7x7 的矩阵。当每一个线程只计算一个结果的时候，我们需要从 A 中读取 7 个数据，从 B 中读取 7 个数据，从 C 中读取 1 个数据，然后写一次 C。这样的话，每个线程需要读取 15 个数据，写一次数据。如果我们每一个线程计算 4 个结果，那么我们需要从 A 中读取 28 个数据，从 B 中读取 7 个数据，从 C 中读取 4 个数据，然后写 4 次 C。这样的话，每个线程需要读取 32 个数据，写 4 次数据。计算每个线程的平方结果比计算结果的列更有效，因为这样我们可以共享更多的输入。 | |||
There was a problem hiding this comment.
这里 A 就是读取 14 个数据，往 C 里面写 4 个数据，就需要读 A 的俩行，B 的俩行，每行 7 个数据
| @@ -6,15 +6,21 @@ | |||
|
|
|||
| 在介绍二维 Thread Tile 之前，我们先来回顾一下一维 Thread Tile 的优化方法。在初级系列中，我们使用了一维线程块来优化矩阵乘法的性能，我们将矩阵乘法的计算任务分配给了一维线程块，每个线程块负责计算一个小的矩阵块。这样做的好处是可以充分利用共享内存，减少全局内存的访问次数，从而提高矩阵乘法的性能。 | |||
|
|
|||
| 我们在每个线程中计算了一维的矩阵块。想要继续优化这个 Kernel 的性能，我们可以使用二维线程块来计算二维的矩阵块。 | |||
| 还记得一维 Thread Tile 中的例子吗？如果输入的 A 和 B 都是 8x8 的矩阵： | |||
There was a problem hiding this comment.
已经修改，可能一开始考虑到 7x7 不会被整除
|
如果内核函数对寄存器或共享内存的使用量较高，使用 这个后续感觉可以在新手村系列加一篇文章单独介绍一下，做一点简单的小实验对比下它的效果，你要是有时间欢迎提 PR 😄 |
muyuuuu
force-pushed
the
gemm_tiled_2d
branch
from
e272187 to
7504a19
Compare
muyuuuu
force-pushed
the
gemm_tiled_2d
branch
from
7504a19 to
66b028d
Compare
查了下文档： 有个疑问，运行 kernel 的时候会指定 block size，也就是每个 block 的线程数是确定的，为啥还有 |
在编译阶段，编译器并不知道你在运行时会用什么线程数。因此，它需要依赖类似 如果编译器认为你可能会用到 1024 个线程，而实际运行时只用到 256 个线程，寄存器可能被过度保守地分配，降低了效率。 max_threads_per_block 还提供了一层安全检查，如果运行时的 threads 超过 max_threads_per_block，内核调用会失败。 两者的目标不同： |
__launch_bounds__((BM * BN) / (TM * TN), 1)，好像没有相关的解释。才疏学浅，我就不乱讲了，大佬可否补充一下？